Drove run_mori.py to a correct run on 8x MI355X (on-node via salloc+srun):
dispatch+combine numerically correct (combine within tol, max_rel ~2e-3),
~85us round-trip at the decode shape. The first runs surfaced four issues,
all fixed and re-validated:

- RDMA MR ceiling: MoRI registers the WHOLE symmetric heap as one RDMA MR at
  init (even single-node; no disable-RDMA knob). The ionic_rdma NICs cap GPU
  MRs at ~4 GiB — a 6 GiB heap fails (RegisterRdmaMemoryRegion errno 22), 2 GiB
  registers. Hold heap at MORI_SHMEM_HEAP_SIZE=2G (override CX_MORI_HEAP_SIZE).
- Buffer sizing: max_num_inp_token_per_rank 4096 -> max(512, n) so the buffers
  fit the 2 GiB heap (4096 was inherited from the reference test).
- Correctness shape: combine returns the full max-token buffer; compare only
  combined[:n] against expected.
- recv count: read total_recv BEFORE combine (combine resets recv_num, which
  made recv_nonzero a false negative).
- Teardown: MoRI's shmem teardown asserts (CheckStatusValid -> SIGABRT) when the
  op is destroyed after shmem_finalize(); hard-exit after writing results.

Docs (README/plan/CONTAINERS) updated from "scaffolded" to validated, with the
fabric constraints recorded.

…t cancel each other

Cleanup (goal): launchers/ now holds ONLY launch_<sku>.sh. Moved the in-container runtime
(common.sh, run_in_container.sh) -> runtime/ and the dev/orchestration helpers (_gha_*,
_keep_newest.py, _v3/_v4/_mi355x/_gb300/_routing/_repro/_sensitivity/_validate*) -> tools/.
Updated launch scripts' source + run_in_container paths ($HERE/../runtime/common.sh;
runtime/run_in_container.sh), run_in_container's CWD-relative common.sh source, + README/
CONTAINERS refs. No functional change to the launch path — validated below via GHA.

DeepEP fp8 dispatch gains scale-layout recipe variants (same kernel, different cast):
fp8=per-block-128 (default), fp8-pertoken (one scale/token), fp8-directcast (unscaled — zero
scale-transport, the recipe MoRI PR311 replaced). Selected via dispatch_dtype; recorded in
backend_provenance.scale_layout so each is a distinct operating point. Closes goal P1
'Direct-cast FP8 dispatch' + 'Per-token/Per-block scale layout variants' + informs 'Scale
transport overhead'. capability/schema/workflow dtype enums extended.

…e backend tracebacks

- ep_harness --dispatch-dtype choices += fp8-pertoken/fp8-directcast (argparse rejected them
  before the adapter ran — the 6 fp8-recipe runs failed here, not in the cast code).
- allreduce_fw_bench status=valid iff nccl baseline produced bw>0 (framework custom kernels
  not in every image are recorded in frameworks_available, not a failure).
- run_ep.py prints any backend exception's full traceback to STDOUT (torchrun summarizes stderr
  only, hiding new-adapter errors like flashinfer's in CI).

The MoeAlltoAll Mapping requires world_size==tp*pp*cp and realizes MoE-EP as a view over the
TP dim, so pure EP across all ranks = tp_size=world_size, moe_ep_size=world_size (was tp_size=1
-> 'world_size must equal tp*pp*cp, 8!=1*1*1'). Confirms the failure was a Mapping-arg bug, NOT
the MNNVL/pidfd hardware blocker — re-smoking to see if FlashInfer EP runs on x86_64.

… worked)

Mapping now constructs (tp_size=world_size); next: MoeAlltoAll asserts hidden_size required.
Pass hidden_size=args.hidden. Confirms steady progress past construction — not a hardware block.

…l unblock

The MoeAlltoAll workspace bootstrapped its cross-rank comm via MPI (MpiComm().Split), failing
under torchrun (no mpi4py/MPI launch) — that was the actual blocker, NOT CAP_SYS_PTRACE. FlashInfer
ships TorchDistBackend + MnnvlConfig; register it via MnnvlMemory.set_comm_from_config(mapping,
config) before constructing MoeAlltoAll so the symmetric workspace uses the torch.distributed NCCL
group torchrun already set up. Pass mnnvl_config to MoeAlltoAll too.

…pace=False

dispatch + MNNVL workspace now WORK (torch-dist fix proved it's not hardware-gated). combine
failed: payload_in_workspace=True demands the payload at the exact workspace pointer (RuntimeError,
not TypeError, so _call_variants didn't fall through). Clone the recv to a fresh tensor + pass
payload_in_workspace=False so the kernel copies our identity-expert output in itself.

MoeAlltoAll is idle->dispatched->idle and asserts 'dispatch called twice without combine'. The
harness times dispatch in isolation (loops it). Fix: combine_needs_redispatch=True (combine-timing
runs untimed dispatch+stage via pre= before each combine sample) + _reset_moe_fsm() at dispatch
start (so the looped dispatch timing + paired roundtrip stay legal). dispatch+workspace already
proven working on x86_64 via torch-dist — this is the last structural gate before correctness.

… FSM)

Looping isolated dispatch corrupts MoeAlltoAll's symmetric workspace (CUDA launch failure) — the
FSM requires PAIRED dispatch+combine. Add a roundtrip_only backend flag: run_sweep then times ONLY
the paired roundtrip (dispatch->stage->combine, each iter cycling the FSM idle->dispatched->idle
cleanly) and mirrors it into dispatch/combine for schema/plot. FlashInfer EP set roundtrip_only=True.
The roundtrip is goal P0's headline metric, so this is the correct measurement here. (Proven: FlashInfer
EP RUNS on x86_64 via torch-dist MNNVL — not the hardware block the early probe assumed.)

…ame [ep,maxT,H] shape)

FlashInfer source: dispatch->recv [ep_size,max_tokens,hidden]; combine wants payload of the SAME
shape. recv[0] IS the identity-expert output -> hand it straight to combine, no clone (the clone
broke the workspace-view layout -> async CUDA corruption). payload_in_workspace=False so the kernel
stages it. + a one-time rank0 shape log. (FlashInfer EP construct+dispatch+MNNVL-workspace proven
working in the GHA cap-enabled container; this is the combine-contract fix.)

FlashInfer EP RUNS on H100 (8 pts, disp_p99=120.7us, conformant, anomaly-free) — only
semantic_correctness failed. combine takes no gate weights + reduces recv [ep_size,max_tokens,H]
over the per-RANK axis, so identity round-trip = x * distinct_ranks_per_token (like DeepEP normal),
not x*topk. Compute it vectorized from the routing trace; default CX_FLASHINFER_ROUTING_FACTOR=ranks.

…import it

cx_build_uccl clones uccl + copies ep/deep_ep_wrapper/deep_ep -> a uccl_deepep package (relative
imports, no deep_ep shadow). ep_uccl.py imports uccl_deepep.Buffer (group ctor, matching its
Buffer(self.group,...) calls + DeepEP's API) with low-level uccl.ep fallback. The earlier UCCL run
failed because uccl.ep.Buffer is (rank,num_ranks) not (group) — the wrapper is the DeepEP-compatible layer.

…v0.1.1)

Wrapper construct worked (group Buffer), but its dispatch hit get_rdma_buffer signature mismatch:
the main-branch wrapper vs the pip wheel 0.1.1. Clone uccl at the tag matching the installed
version so wrapper + uccl.ep C-extension agree.

…already written)

UCCL EP RUNS + PASSES (H100 smoke: status=valid, correct=True x8, disp_p50=146us, comb_p50=105us)
but SIGSEGVs in symmetric-memory teardown after the JSON is written. Like ep_mori, os._exit(rc) past
the crashy cleanup so a valid result isn't marked failed. UCCL EP now a working backend (v0.1.1
wrapper + cu12 shim).

…+ nvfp4)

Wire quantized dispatch into the FlashInfer EP adapter (MoeAlltoAll, the TRT-LLM
throughput-backend one-sided A2A). The A2A is a dtype-agnostic byte-mover taking
input_payloads as a list, so a quantized dispatch moves [q, scale_factor] and dequants
in stage(). Adds:
 - fp8/fp8-pertoken/fp8-directcast (e4m3, same convention as ep_deepep)
 - mxfp8 (e4m3 + e8m0 block-32, device dequant verified == flashinfer host dequant)
 - nvfp4 (e2m1 + e4m3 block-16, fp4_quantize + e2m1_and_ufp8sf_scale_to_float)
Dequant is cached on the problem (deterministic recv) so the roundtrip measures comm
only, consistent with DeepEP's untimed-stage timing boundary. mxfp4 excluded (its SF is
tile-padded, not per-token-movable through the A2A). Validated end-to-end EP2 on B300:
all correct=True, comm-only latency. Records trtllm lineage in provenance (MoeAlltoAll
lives in flashinfer.comm.trtllm_moe_alltoall) + scale_layout/quant_kind.

…reduce_fusion)

Replace the guessed flashinfer custom-AR entrypoints with the pinned real contract
(flashinfer 0.6.8.post1): trtllm_allreduce_fusion with pattern_code=AllReduceFusionPattern.
kAllReduce (pure AR, no fusion) + use_oneshot True/False = one-shot vs two-shot, over the IPC
workspace from trtllm_create_ipc_workspace_for_all_reduce_fusion. Adds a _SkipSize path so sizes
a kernel can't shape (sub-hidden; two-shot needs token_num>tp_size) are recorded as skipped rows
without failing the impl. Validated EP2 on B300: nccl 56.5, flashinfer-oneshot 59.9 (beats nccl
in the latency regime), flashinfer-twoshot 36.3 GB/s, all correct. Covers goal's FlashInfer /
one-shot / two-shot all-reduce.

nvfp4 (e2m1 FP4) is a Blackwell-native format — FlashInfer's fp4 quantize/dequantize does not
round-trip on Hopper sm90 (validated: correct=True on B300, correct=False on H100). Add
ARCH_ONLY_DTYPES + _sku_arch() so capability.resolve rejects nvfp4 on Hopper cleanly (mxfp8/e4m3
unaffected). Refresh docs/gated.md: FlashInfer EP is DONE on H100+B300 (the prior x86_64-blocked
framing was wrong — only the H200 runner denies CAP_SYS_PTRACE); MXFP8+NVFP4 dispatch done (mxfp4
gated on its tile-padded SF); quant-combine OUTPUT gated (no output_dtype in flashinfer 0.6.8.post1);
FlashInfer one-shot/two-shot all-reduce done.

…d.md refresh

plot_ep.py: add load_allreduce_fw_series() — render family=allreduce-fw (flashinfer one-shot/
two-shot vs nccl) in the existing All-reduce tab (op=all_reduce, one line per impl, per-SKU-family
colors). Purely additive; no JS change. docs/gated.md already refreshed in the prior commit covers
FlashInfer EP done on H100+B300, mxfp8/nvfp4 dispatch, quant-combine-output gated, framework AR.

…e / all-gather / framework AR)

Add a methodology.md section mapping the non-EP collective families to their serving patterns:
TP-activation all-reduce (NCCL ring vs FlashInfer one-shot/two-shot crossover), and the DP-attention
-> TP-MoE all-gather handoff (the standardized byte sweep spans the [total_tokens,hidden] handoff
payload sizes). Documents the SGLang DP-attention use case honestly (size coverage present; named
per-model shapes = further lift).

Add deepep-hybrid as a 5th NVIDIA EP backend. The hybrid-ep branch (NVIDIA's TMA + warp-pipeline
impl, deep_ep.HybridEPBuffer) is built from source by cx_build_deepep_hybrid with 3 container fixes
(CUDA-13 cccl include via CPATH; unversioned libnvshmem_host.so symlink; NVSHMEM_DIR). ep_deepep_
hybrid.py wraps HybridEPBuffer (distinct API: construct with hidden/max_tokens/local_experts;
dispatch(hidden,topk_idx=,topk_weights=,num_of_experts=) -> (recv,_,_,handle); combine(recv,handle=)).
Intranode NVLink path (<=8 ranks, one NVLink domain). Identity-expert combine uses the per-rank-sum
'ranks' factor (verified: 8-rank topk=8 relerr 4.28 == E[distinct ranks] 5.26). Validated end-to-end
EP2 + EP8 on B300: correct=True, status=valid. Wired capability (Blackwell+Hopper, bf16 normal),
run_ep --backend, run_in_container CX_BENCH case, schema enum, workflow option. fp8 (use_fp8) +
internode NVLink<->RDMA forwarding are further lift (multi-node).

…cache/rl-mesh/allreduce-fw)

launch_mi355x-amds.sh previously forced CX_BENCH=mori for anything except mori/nccl, blocking the
AMD-capable collective families. Extend the allow-list to mori|nccl|kv-cache|rl-mesh|allreduce-fw
(all run on ROCm/CDNA4: HIP memcpy for kv-cache, torch.distributed->RCCL for rl-mesh/allreduce-fw;
allreduce-fw's flashinfer one/two-shot self-skip on the ROCm image, leaving a valid RCCL baseline).
NVIDIA-only benches still fall back to mori (and capability rejects them on amd anyway).

…a2a_combine

Confirmed from the flashinfer main-branch source that moe_a2a_combine/MoeAlltoAll.combine gained
output_dtype + output_scales + output_scalar_scale (bf16/fp8_e4m3fn/packed-fp4) in a newer release
(PR3376/3643) — absent in the bundled 0.6.8.post1. Wire a quantized COMBINE OUTPUT path (fp8 e4m3):
 - run_in_container cx_build_flashinfer_latest: pip -U flashinfer when CX_COMBINE_DTYPE!=bf16 (asserts
   the upgraded combine has output_dtype); dispatch path stays on the bundled version.
 - ep_flashinfer combine(): single-shot quant-combine (output_dtype=fp8 + output_scales[T,1] per-token,
   CX_QC_SCALE override), dequant cached/untimed for correctness (the fp8 reduction is timed).
 - capability flashinfer combine_dtypes={bf16,fp8}; harness/run_ep gate already reads SUPPORTED_COMBINE_*.
 - workflow: input-cap-safe 'flashinfer-combine-fp8' benchmark choice -> CX_BENCH=flashinfer +
   CX_COMBINE_DTYPE=fp8 (capability-validate mapped). nvfp4/mxfp8 combine reserved until fp8 validates.

…ersions

The first combine-fp8 GHA run upgraded flashinfer-python 0.6.8.post1 -> 0.6.13 but left
flashinfer-cubin at 0.6.8.post1, so import flashinfer raised a version-mismatch (the combine
output_dtype IS in 0.6.13). Upgrade flashinfer-python + flashinfer-cubin + flashinfer-jit-cache
together and export FLASHINFER_DISABLE_VERSION_CHECK=1 as a fallback.

…ter)

MoRI runs on the shared MI355X timed out at the default 900s per-phase guard (cold enroot import +
node contention + MoRI's slowness at larger T). Raise CX_RUN_TIMEOUT to 1800s in launch_mi355x-amds.sh
(fits the 60-min salloc). Pairs with reduced --timing on re-dispatch.

…_dtype (not in 0.6.13 PyPI)

The combine output_dtype is on flashinfer main but not the latest PyPI release (0.6.13), so pip -U
was insufficient (asserted 'combine still has no output_dtype'). Install from the nightly wheel
index https://flashinfer.ai/whl/nightly/ (flashinfer-python --no-deps + cubin + cu130 jit-cache),
which is built from main and has the quantized-combine output path.

…CuTe DSL coupling)

The nightly (main) flashinfer's CuTe-DSL kernels import newer cutlass.cute symbols
(OperandMajorMode) than the container's bundled nvidia-cutlass-dsl — import failed. Upgrade
nvidia-cutlass-dsl alongside the nightly flashinfer-python/cubin/jit-cache.

…ance

The flashinfer nightly upgrade happens AFTER env_capture, so the upgraded versions live nowhere
else. cx_build_flashinfer_latest now captures CX_FLASHINFER_STACK (flashinfer-python/cubin/jit-cache
+ nvidia-cutlass-dsl + torch versions) — logged to the GHA log AND read into ep_flashinfer's
backend_provenance (flashinfer_stack). Reproducibility for the quant-combine runs (which depend on a
specific newer flashinfer+cutlass-dsl set).